@llvm/pr-subscribers-clang

Author: Denis.G (DenisGZM)

diff --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp
index 150b2879fc94f..33b9f63bcfa08 100644
--- a/clang/lib/Parse/ParseStmt.cpp
+++ b/clang/lib/Parse/ParseStmt.cpp
@@ -296,6 +296,11 @@ StmtResult Parser::ParseStatementOrDeclarationAfterAttributes(
     goto Retry;
   }
 
+  case tok::kw_alignas: {
+    ParseAlignmentSpecifier(CXX11Attrs);
+    goto Retry;
+  }
+
   case tok::kw_template: {
     SourceLocation DeclEnd;
     ParseTemplateDeclarationOrSpecialization(DeclaratorContext::Block, DeclEnd,
diff --git a/clang/test/SemaCUDA/cuda-attr-order.cu b/clang/test/SemaCUDA/cuda-attr-order.cu
new file mode 100644
index 0000000000000..d3bf5b014d1c6
--- /dev/null
+++ b/clang/test/SemaCUDA/cuda-attr-order.cu
@@ -0,0 +1,15 @@
+// Verify that we can parse a simple CUDA file with different attributes order.
+// RUN: %clang_cc1 "-triple" "nvptx-nvidia-cuda"  -fsyntax-only -verify %s
+// expected-no-diagnostics
+#include "Inputs/cuda.h"
+
+struct alignas(16) float4 {
+    float x, y, z, w;
+};
+
+__attribute__((device)) float func() {
+    __shared__ alignas(alignof(float4)) float As[4][4];  // Both combinations
+    alignas(alignof(float4)) __shared__  float Bs[4][4]; // must be legal
+
+    return As[0][0] + Bs[0][0];
+}

CC @erichkeane in case there’s a specific reason I’m not aware of as to why we currently don’t allow this.

At a glance this does seem like the right place to do this, but this is still missing a release note.

It seems like GCC allows e.g. __attribute__(()) alignas(16) int x in any case, so I don’t see why we shouldn’t allow this too. Can you also add some tests that use __attribute__(()) directly and which aren’t CUDA-specific?

Oh, and can you add solmething like this as a test as well:

struct S { __attribute__((deprecated)) alignas(16) int x; };

Actually this test doesn't work with this patch...

In this case all attributes are processed in ParseDeclarationSpecifiers, which in my first view was the right place to fix, but has way more complicated logic and easy to break diagnostics.

In ParseDeclarationSpecifiers we parse kw__attributes and other CXX11 Attributes and set bool AttrsLastTime = true to check that last parsed piece was attr. Later this block prohibit attributes with AttrsLastTime = false,

ParseDecl.cpp

    DoneWithDeclSpec:
      if (!AttrsLastTime)
        ProhibitAttributes(attrs);

And AttrsLastTime is always false in declarations of the form: <attributes> <type> <identifier> , because last token we parse is type

Another approach i tried is to add processing alignas-cxx11 just like it is done for C: kw__Alignas and kw_alignas (c23).
Well, it do the parsing but later it skips CXX11 attributes when correcting declaration type (assumed that attributes must have been processed before)

Hmm, @erichkeane probably knows where this needs to be parsed then; I might take another look at this myself later (because I’m not sure either off the top of my head), but I’m rather busy today unfortunately...

I added parsing all attributes in ParseCXXClassMemberDeclaration before calling ParseDeclarationSpecifiers and it seems to solve problem, but it also changes annotation ranges for struct and class members

Reworked approach for parsing.

We needed to support arbitrary attribute parsing rather than just alignas for CXX. So please check new commit

✅ With the latest revision this PR passed the C/C++ code formatter.